[microNPU] Add MergeConstants pass #12029

NicolaLancellotti · 2022-07-07T10:30:03Z

This pr adds a MergeConstants pass that looks for the constants used by each compute operator and merges them into a single buffer.

@manupa-arm @ekalda @lhutton1

cc @Mousius

Change-Id: I1ff51d8147fba8c66d442a370b9f058e9b2758d8

Change-Id: I29f68f83a73fa00ca34ed0ab2321c53c6b761137

manupak

Thanks @NicolaLancellotti !

I have a done a round of reviews. There are few things we should improve in the PR. See if you agree.

Also please consider adding docs for non-trivial functions.

We need to align with coding style of TVM : https://tvm.apache.org/docs/contribute/code_guide.html
(I have put a few specific comments to each type)

manupak · 2022-07-07T10:58:27Z

python/tvm/relay/backend/contrib/ethosu/tir/compiler.py

@@ -90,6 +90,8 @@ def lower_ethosu(sch, args, const_dict, name="main"):
        mod = tvm.tir.transform.RemoveNoOp()(mod)
        mod, const_dict = ethosu_passes.EncodeConstants(const_dict)(mod)
        mod = ethosu_passes.HoistAllocates()(mod)
+        if not util.is_striping_enabled():


Please leave a comment that MergeConstant pass currently does not support striped schedules and requires further investigation.

manupak · 2022-07-07T10:59:24Z

python/tvm/relay/backend/contrib/ethosu/tir/passes.py

+        new_const_dict = {}
+        for param in const_dict.keys():
+            new_const_dict[tvm.tir.IntImm("int64", param)] = tvm.nd.array(const_dict[param])
+        mod["main"] = mod["main"].with_attr("ethos-u.const-dict", new_const_dict)


nit : lets stick to const_dict

manupak · 2022-07-07T11:02:09Z

src/tir/contrib/ethosu/passes.cc

+    analyze = false;
+    new_body = rewrite_prim_func_body(new_body);
+    std::set<ObjectRef> params_to_delete{};
+    auto new_buffer_map{make_new_buffer_map(main_func->buffer_map, &params_to_delete)};


lets not use auto here.
it is fine to use auto when you can understand the type from just looking at the line. In this case, one need to look into the function.

Quote from from the following :
"For example, you can assume that the return type of make_unique() is obvious, but the return type of MyWidgetFactory() probably isn't."

Reference : https://google.github.io/styleguide/cppguide.html#Type_deduction

(Applies to all the lines)

manupak · 2022-07-07T11:05:29Z

src/tir/contrib/ethosu/passes.cc

+
+    // Rewrite new allocates
+    for (auto it{copy_write_buffers.rbegin()}; it != copy_write_buffers.rend(); ++it) {
+      if (auto buffer_opt = *it) {


Lets use the type here.

manupak · 2022-07-08T09:19:15Z

python/tvm/relay/backend/contrib/ethosu/tir/passes.py

+    Constants written to a buffer with local scope are not merged.
+    """
+
+    def mergeConstantsPass(mod):


Style : this should be lower case. https://peps.python.org/pep-0008/#function-and-variable-names

(Only exception is the Passes because we want them aligned with C++ ones and I think the from API PoV the accessible pass function should be MergeConstants)

manupak · 2022-07-08T09:27:54Z

src/tir/contrib/ethosu/passes.cc

+
+  PrimFunc operator()(PrimFunc main_func, const Map<IntImm, runtime::NDArray>& const_dict) {
+    // Analyze
+    Stmt new_body{this->VisitStmt(main_func->body)};


It is generally a not good practice to mix analysis and transformation in a single pass.
E.g. for analysis pass : (https://github.com/apache/tvm/blob/main/src/tir/usmp/analysis/extract_buffer_info.cc)
However, since this is simpler analysis we can keep the analysis pass internal to this source.

Can we break this out to a new pass (in this source file) that return the analysis information and run the sub-sequent pass to rewrite w/o resorting to a mode variable analyze ?

manupak · 2022-07-08T09:29:22Z

src/tir/contrib/ethosu/passes.cc

+    if (op->size() <= 1) {
+      return StmtExprMutator::VisitStmt_(op);
+    }
+    return analyze ? analyze_seq_stmt(op) : rewrite_seq_stmt(op);


Refer to the above comment about seperation of analysis and transformation passes.

manupak · 2022-07-08T09:32:07Z

src/tir/contrib/ethosu/passes.cc

+    std::map<const VarNode*, Allocate> var_to_allocate{};
+
+    // Rewrite old allocates
+    std::set<ObjectRef> buffer_vars{get_vars_for_written_copy_buffers()};


Do we need a ordered set here ? (Noting that ordered_set's access time is higher)

I switched to unordered_set.

manupak · 2022-07-08T09:35:23Z

src/tir/contrib/ethosu/passes.cc

+    for (size_t i{0}; i < seq_stmt.size(); ++i) {
+      Stmt stmt{seq_stmt[i]};
+
+      switch (get_stmt_type(stmt)) {


We need a default here.
Ref : https://google.github.io/styleguide/cppguide.html#Loops_and_Switch_Statements

It is a switch on an enumerated value, and all cases are covered. So we don't need a default, don't we?

manupak · 2022-07-08T09:35:34Z

src/tir/contrib/ethosu/passes.cc

+    for (size_t i = 0; i < seq_stmt.size(); ++i) {
+      Stmt stmt{seq_stmt[i]};
+
+      switch (get_stmt_type(stmt)) {


We need a default here.
Ref : https://google.github.io/styleguide/cppguide.html#Loops_and_Switch_Statements

It is a switch on an enumerated value, and all cases are covered. So we don't need a default, don't we?

Change-Id: Iad59107d5abdec6b079c6fd4ab48c6bffbb5e0bb

NicolaLancellotti · 2022-07-08T16:39:21Z

Thanks @NicolaLancellotti !

I have a done a round of reviews. There are few things we should improve in the PR. See if you agree.

Also please consider adding docs for non-trivial functions.

We need to align with coding style of TVM : https://tvm.apache.org/docs/contribute/code_guide.html (I have put a few specific comments to each type)

Thanks for the review @manupa-arm!

Change-Id: Ie5caf506337de01e169d6f422e4682eefbd93241

manupak

LGTM!

manupak · 2022-07-12T16:23:54Z

Thanks @NicolaLancellotti !

* [microNPU] Add MergeConstants pass Change-Id: I1ff51d8147fba8c66d442a370b9f058e9b2758d8 * Fix errors and warnings Change-Id: I29f68f83a73fa00ca34ed0ab2321c53c6b761137 * Address comments Change-Id: Iad59107d5abdec6b079c6fd4ab48c6bffbb5e0bb * Fix lint error Change-Id: Ie5caf506337de01e169d6f422e4682eefbd93241

Nicola Lancellotti added 2 commits July 7, 2022 08:49

[microNPU] Add MergeConstants pass

cd8b94b

Change-Id: I1ff51d8147fba8c66d442a370b9f058e9b2758d8

Fix errors and warnings

2b8f95a

Change-Id: I29f68f83a73fa00ca34ed0ab2321c53c6b761137

github-actions bot requested a review from Mousius July 7, 2022 13:24

manupak requested changes Jul 8, 2022

View reviewed changes

manupak self-assigned this Jul 8, 2022

Address comments

43bcaad

Change-Id: Iad59107d5abdec6b079c6fd4ab48c6bffbb5e0bb

Fix lint error

5ca38dc

Change-Id: Ie5caf506337de01e169d6f422e4682eefbd93241

manupak approved these changes Jul 12, 2022

View reviewed changes

manupak merged commit fbf80bb into apache:main Jul 12, 2022

AndrewZhaoLuo mentioned this pull request Oct 4, 2022

TVM v0.10.0.rc0 Release Candidate Notes #12979

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

[microNPU] Add MergeConstants pass #12029

[microNPU] Add MergeConstants pass #12029

NicolaLancellotti commented Jul 7, 2022 •

edited by github-actions bot

Loading

manupak left a comment •

edited

Loading

manupak Jul 7, 2022

manupak Jul 7, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 7, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 7, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

manupak Jul 8, 2022

NicolaLancellotti Jul 8, 2022

NicolaLancellotti commented Jul 8, 2022

manupak left a comment

manupak commented Jul 12, 2022

[microNPU] Add MergeConstants pass #12029

[microNPU] Add MergeConstants pass #12029

Conversation

NicolaLancellotti commented Jul 7, 2022 • edited by github-actions bot Loading

manupak left a comment • edited Loading

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

NicolaLancellotti commented Jul 8, 2022

manupak left a comment

Choose a reason for hiding this comment

manupak commented Jul 12, 2022

NicolaLancellotti commented Jul 7, 2022 •

edited by github-actions bot

Loading

manupak left a comment •

edited

Loading